﻿///////////////////////////////////////////////////////////////////////////////
// This implements the separable convolution
// use only the texture memory
///////////////////////////////////////////////////////////////////////////////

#include <math.h>
#include "convolutionTexture.h"
#include "../common/common.h"

#define IMAGE_WIDTH 1024
#define IMAGE_HEIGHT 1024

int main(int argc, char **argv)
{
	float	*h_VerticalKernel,
			*h_HorizontalKernel,
			*h_Input,
			*h_Buffer,
			*h_OutputCPU,
			*h_OutputGPU;

	cudaArray *a_Src;

	cudaChannelFormatDesc floatTex = cudaCreateChannelDesc<float>();

	float	*d_Output;

	// Cấp phát bộ nhớ trên CPU và GPU
	h_VerticalKernel = (float *)malloc(KERNEL_LENGTH * sizeof(float));
	h_HorizontalKernel = (float *)malloc(KERNEL_LENGTH * sizeof(float));
	h_Input     = (float *)malloc(IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float));
	h_Buffer    = (float *)malloc(IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float));
	h_OutputCPU = (float *)malloc(IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float));
	h_OutputGPU = (float *)malloc(IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float));
	
	SAFE_CALL( cudaMallocArray(&a_Src, &floatTex, IMAGE_WIDTH, IMAGE_HEIGHT) );
	SAFE_CALL( cudaMalloc((void **)&d_Output, IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float)) );

	srand(2009);
	for(unsigned int i = 0; i < KERNEL_LENGTH; i++){
		h_VerticalKernel[i] = (float)(rand() % 16);
		h_HorizontalKernel[i] = (float)(rand() % 16);
	}

	for(unsigned int i = 0; i < IMAGE_WIDTH * IMAGE_HEIGHT; i++)
		h_Input[i] = (float)(rand() % 16);

	setConvolutionKernel(h_VerticalKernel, h_HorizontalKernel);
	SAFE_CALL( cudaMemcpyToArray(a_Src, 0, 0, h_Input, 
		IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float), cudaMemcpyHostToDevice) );


	// Gọi hàm convolutionColumnGPU trên GPU
	convolutionColumnGPU(d_Output, a_Src, IMAGE_WIDTH, IMAGE_HEIGHT);

	//While CUDA kernels can't write to textures directly, this copy is inevitable
	printf("Copying convolutionRowGPU() output back to the texture...\n");
	SAFE_CALL( cudaThreadSynchronize() );
	SAFE_CALL( cudaMemcpyToArray(a_Src, 0, 0, d_Output, 
		IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float), cudaMemcpyDeviceToDevice) );
	SAFE_CALL( cudaThreadSynchronize() );

	// Gọi hàm convolutionRowGPU trên GPU
	convolutionRowGPU(d_Output, a_Src, IMAGE_WIDTH, IMAGE_HEIGHT);
	SAFE_CALL( cudaThreadSynchronize() );

	// Chép kết quả về CPU
	SAFE_CALL( cudaMemcpy(h_OutputGPU, d_Output, 
		IMAGE_WIDTH * IMAGE_HEIGHT * sizeof(float), cudaMemcpyDeviceToHost) );

	printf("Kiem tra ket qua...\n");
	printf("...goi ham convolutionColumnCPU()\n");
	convolutionColumnCPU(h_Buffer, h_Input, h_VerticalKernel,
		IMAGE_WIDTH, IMAGE_HEIGHT, KERNEL_RADIUS);

	printf("...goi ham convolutionRowCPU()\n");
	convolutionRowCPU(h_OutputCPU, h_Buffer, h_HorizontalKernel,
		IMAGE_WIDTH, IMAGE_HEIGHT, KERNEL_RADIUS);

	double delta = 0;
	double sum = 0;
	for(unsigned int i = 0; i < IMAGE_WIDTH * IMAGE_HEIGHT; i++){
		if (i < 10)
		{
			printf("%.2f %.2f\r\n", h_OutputCPU[i], h_OutputGPU[i]);
		}
		sum += h_OutputCPU[i] * h_OutputCPU[i];
		delta += (h_OutputGPU[i] - h_OutputCPU[i]) * (h_OutputGPU[i] - h_OutputCPU[i]);
	}
	double L2norm = sqrt(delta / sum);
	printf("Relative L2 norm: %E\n", L2norm);
	printf((L2norm < 1e-6) ? "PASSED\n" : "FAILED\n");


	// Giải phóng bộ nhớ
	SAFE_CALL( cudaFree(d_Output) );
	SAFE_CALL( cudaFreeArray(a_Src) );
	free(h_OutputGPU);
	free(h_Buffer);
	free(h_Input);
	free(h_VerticalKernel);
	free(h_HorizontalKernel);

	cudaThreadExit();
}
